\chapter{Introduction au domaine}

Ces travaux sont à l'intersection de deux domaines:
\begin{itemize}
 \item La programmation massivement multiprocesseurs avec la technologie CUDA
 \item Le traitement par matrice de convolution.
\end{itemize}

\section{CUDA}

CUDA pour Compute Unified Device Architecture, est une technologie développée par {\Large \textsc{nVidia}}. Elle est basée sur une architecture de type GPU \footnote{GPU : Graphics Processing Unit} massivement multiprocesseurs ainsi que sur une bibliothèque permettant son utilisation. Cette technologie est supportée par les derniers modèles de carte graphique {\Large \textsc{nVidia}}.

\subsection{Les origines de CUDA}

Dans un premier temps, l'accroissement de la puissance de calcul des processeurs consistait à augmenter le nombre de transistors et la fréquence, réduisant régulièrement la taille de gravure des processeurs. À partir de 2004, cette solution posa des problèmes de dissipation thermique. Les constructeurs de processeurs s'orientèrent donc vers l'augmentation du nombre de processeurs par puce. En 2006, avec la GeForce 8800 GTX, {\Large \textsc{nVidia}} sort la première carte graphique possédant un GPU avec une architecture CUDA.

\subsection{Architecture}

\subsubsection{Hiérarchisation des threads}

\paragraph{}Pour effectuer des calculs sur les GPU CUDA nous emploierons des kernels. Un kernel est une portion de code à exécuter en parallèle sur la carte graphique. Chaque instance du kernel est donc un thread. Les threads sont regroupés par blocs. Chaque bloc est indépendant ce qui implique que les threads ne peuvent pas communiquer directement avec les threads des autres blocs. Cette organisation des threads est liée à l'organisation du GPU. En effet, les GPU CUDA sont composés de multiprocesseurs, chaque multiprocesseur étant composé de processeurs (voir figure \ref{fig:archi}).

\newpage
\begin{figure}[H]
  \centering
  \includegraphics[height=9cm,width=14cm]{images/machinkicontientdestrucsenglobedebidul.png}
  \caption{Schéma de l'architecture CUDA.}
  \label{fig:archi}
\end{figure}

\paragraph{}A l'exécution, les threads d'un même bloc sont exécutés par warps, un warp étant un ensemble de $32$ threads sur les architectures actuelles. Les threads d'un warp sont exécutés de manière synchrone, instruction par instruction.\\

\begin{figure}[H]
  \centering
  \includegraphics[height=7cm,width=12cm]{images/grille-bloc.png}
  \caption{Hiérarchisation des threads avec une grille en 2 dimensions et des blocs en 3 dimensions.}
\end{figure}

\subsubsection{Répartition des mémoires}

Une carte graphique basée sur l'architecture CUDA possède plusieurs niveaux de mémoire, présentés ici du niveau le plus éloigné du GPU au niveau le plus proche.

\paragraph{}Tout d'abord la mémoire globale est accessible en lecture et en écriture par le GPU, ainsi que par le CPU. Elle est de grande taille, de quelques $Mo$ à quelques $Go$, et est plus rapide que la mémoire centrale du CPU; mais reste la mémoire la plus lente de la carte graphique.

\paragraph{}Une partie de la mémoire globale peut être reservée pour un usage ``texture''. Définir une mémoire en tant que texture permet d'accélérer les temps d'accès en lecture par rapport à la mémoire globale. Cette mémoire utilise des mécanismes propres aux cartes graphiques pour permettre un accès rapide à la totalité de la texture, quelque soit l'unité de calcul appelante. Cependant, la taille d'allocation dans cette mémoire est limitée.

\paragraph{}Ensuite la mémoire constante, stockée dans une zone réservée de la mémoire globale, est accessible en lecture et en écriture par le CPU et en lecture uniquement par le GPU. Elle offre des temps d'accès plus rapides pour les threads que la mémoire globale car elle peut exploiter un mécanisme de cache, et ses accès peuvent être effectués en parallèle.

\paragraph{}La mémoire shared\footnote{Mémoire partagée qui est accessible par plusieurs unités de calculs.}, embarquée dans les Multi-Processeurs du GPU, est accessible en écriture et lecture depuis l'ensemble des Processeurs du Multi-Processeur. Chaque thread d'un bloc ne peut accéder qu'à la mémoire shared de son bloc. Elle est divisée en banques. Chaque thread peut accéder simultanément à des données dans des banques différentes (par un mécanisme de broadcast). Si deux threads tentent d'accéder simultanément à des données différentes d'une même banque, il y a un conflit et les accès sont sérialisés. Cette mémoire offre des temps d'accès proche de ceux des registres, donc très rapides, mais est limitée à quelques kilo-octets.

\paragraph{}Enfin, chaque processeur possède des registres qui sont les mémoires les plus rapides disponibles.
\begin{figure}[H]
  \centering
  \includegraphics[height=8cm,width=8cm]{images/memoire.png}
\caption{Répartition de la mémoire sur une carte graphique reposant sur une architecture CUDA}
\end{figure}

\subsection{Les outils fournis pour utiliser CUDA}

CUDA fournit des outils pour pouvoir développer des applications qui tirent partie de l'architecture CUDA. Ces outils sont :
\begin{itemize}
\item un pilote
\item un SDK\footnote{Software Development Kit}
\item des bibliothèques (par exemple CUBLAS, CUFFT)
\end{itemize}

Il est possible de développer des applications qui reposent sur différents niveaux, les bibliothèques pour le plus haut niveau, et le pilote directement pour le plus bas niveau. Une application programmée à plus bas niveau peut comporter des optimisations plus spécifiques, mais perd en généralisation.\\

\begin{figure}[H]
  \centering
  \includegraphics[height=5cm,width=5cm]{images/niv-api.png}
\caption{Les différents niveaux}
\end{figure}

\subsubsection{Bibliothèques - Qualifieurs de fonctions}

\paragraph{}Ces bibliothèques présentent une extension des langages C, C++ et Fortran. Elles sont accompagnées d'un compilateur, $nvcc$ qui permet de compiler les fichiers les utilisant. Dans les bibliothèques CUDA le CPU est appelé $host$ et le GPU est appelé $device$. Les kernels sont les fonctions devant être exécutées parallèlement sur le GPU.

\paragraph{}Les fonctions doivent avoir une spécification supplémentaire indiquant par quel périphérique elles peuvent être appelées et par quel périphérique elles doivent être exécutées.

\begin{itemize}
\item Les fonctions appelées et exécutées par le CPU sont déclarées avec le mot clé \texttt{\_\_host\_\_}. Par défaut les fonctions sont déclarées en $host$ et donc ce mot clé n'est pas obligatoire.
\item Les fonctions appelées par le CPU et exécutées par le GPU doivent être déclarées avec le mot clé \texttt{\_\_global\_\_}.
\item Les fonctions appelées et exécutées par le GPU doivent être déclarées avec le mot clé \texttt{\_\_device\_\_}. Une fonction peut être en même temps $device$ et $host$.
\end{itemize}

\paragraph{}Lors de l'appel d'une fonction globale, il est nécessaire de préciser la taille de la grille et des blocs. Ces dimensions sont précisées entre triple chevrons, placés juste après le nom de la fonction. Chaque dimension doit être précisée à l'aide d'une instance de la structure dim3.\\

\begin{Verbatim}
 __global__ void mafonction(int a, int b) {
 [...]
 }
 [...]
 dim3 grilleD(3,3,1);
 dim3 blocD(3,3,3);
 mafonctionglobale < < < grilleD, blocD > > > (a, b);
\end{Verbatim}


\subsubsection{Bibliothèque - Qualifieurs de variables}

Les variables peuvent être stockées dans les différentes mémoires de la carte graphique. Pour indiquer explicitement où stocker une variable il suffit lors de la déclaration d'ajouter avant le type l'un des qualifieurs suivants :
\begin{itemize}
\item \texttt{\_\_device\_\_} : variable stockée dans la mémoire globale,
\item \texttt{\_\_constant\_\_} : constante stockée dans la mémoire constante,
\item \texttt{\_\_shared\_\_} : variable stockée dans la mémoire shared.
\end{itemize}

\subsubsection{Bibliothèque - autre}

La bibliothèque propose de nombreuses autres fonctions pour synchroniser les threads d'un bloc ou encore accèder à des fonctions mathématiques plus complexes.

\section{Les matrices de convolution}

La plupart des filtres de traitement des images utilisent des matrices de convolution. Elles permettent, entre autres, de faire des flous, de la détection de contours, des translations. Mais l'utilité des matrices de convolutions ne s'arrête pas au domaine du traitement des images, de nombreux problèmes peuvent être traités par des matrices de convolution.

\subsection{Définition}

\paragraph{}Le principe d'une matrice de convolution est simple. Il s'agit de définir une matrice dite ``de convolution'' et de l'appliquer sur la matrice de données en calculant successivement (pour chaque pixels/donnée) la valeur résultante.
\paragraph{}Généralement, on divise la valeur résultante, la somme des produits voisins/valeur matrice étant généralement bien supérieure au 255 maximal d'une composante d'un pixel, pour la ramener dans des bornes acceptables.

\begin{figure}[H]
  \centering
  \includegraphics[]{images/convolution-calculate.png}
\caption{Exemple simple de convolution}
\end{figure}

L'avantage de cette méthode est qu'il est possible de diviser nos
données et d'effectuer les calculs en parallèle. En effet, il n'y a
pas de lien direct entre chaque donnée, il faut uniquement gérer les
bords lors du découpage.

\newpage
\subsection{Les algorithmes assimilés}

Les algorithmes assimilés à un calcul de matrice de convolution
nécessitent une étape supplémentaire, elle consiste à calculer la
matrice de convolution à appliquer sur la matrice en entrée pour avoir
le résultat escompté. De cette manière, il est possible de résoudre
les problèmes suivants en les résumant à l'application d'une
convolution :
\begin{itemize}
\item Le calcul d'écoulement de fluide,
\item Le calcul de diffusion de chaleur,
\item La détection de contours dans une image (filtre de Sobel).\\
\end{itemize}

Ce qui donnerait pour un écoulement des fluides une décomposition de
l'algorithme de la manière suivante :\\

\begin{Verbatim}
 void calculeEcoulementFluide(int** entree, 
                              int viscosite, 
                              int advection) {

   int** matConvolutionFluide = CalculerMatConvolutionFluide(viscosite, 
                                                           advection);

   [...]

   ConvolEngine engine(matConvolutionFluide);
   engine.convol(entree);

   [...]

 }
\end{Verbatim}

\paragraph{}Il parait donc pertinent d'appliquer des traitements par matrice de convolution de manière massivement parallèle sur architecture GPU.
